perf(gemma4): dual-GPU FP8 resident without host OOM + peer mix - #154
Conversation
Warm decode: MoE ~95% of layer time (attn~4% mlp~1%). Microbench: GemmEx BF16 gate_up~30us down~17us; hipBLASLt FP8 heuristic unsupported on gfx1201. Upstream check: mudler#154 open; decode-graph is CUDA-only (ROCm no graph capture). FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
…re-push sandbox (#159) Two guards on main were RED and between them blocked every open external contributor PR (#127, #154, #155) and every push. Both premises were verified in the tree before changing anything. 1) check-device-leakage: src/vllm/v1/worker/gpu/runner.cpp named vt::DeviceType::kCUDA in the device-agnostic shared layer (DSR bucket 'kcuda' 1 > baseline 0). It came in with the QueueSupportsAsyncInputCombine rescope during the PR #140 fix round - ours, not the contributors'; richiejp reported it in #127's honest gaps. Fixed the way the guard's own message prescribes, mirroring the SupportsAuxStream precedent: ask the backend, not the device. New vt::Backend::SupportsAsyncSampledTokenReadback() (base false) answers whether the host may validly read the sampled token id back between steps; CPU overrides true (host and device memory are one allocation) and CUDA overrides true (the id is device-mirrored). The runner asks vt::TryGetBackend(queue.device.type), whose nullptr for a device absent from the build also subsumes the old #ifdef VLLM_CPP_CUDA guard. SEMANTICS UNCHANGED: CPU async-ON, CUDA async-ON, discrete non-CUDA (ROCm gfx1201) async-OFF - the "!"-token hazard stays closed. 2) .githooks/pre-push ran check-policy.py inside a PARTIAL export (README.md docs scripts .agents), but policy_contract.py:428 asserts AGENTS.md is a non-symlink regular file and resolves its Markdown links against that sandbox. AGENTS.md and its .env.example link were both missing, so the hook failed closed on content that is fine in the real tree - every push refused. EXPORT_PATHS is now a superset of what the CHECKERS read. Gates: check-device-leakage RED->GREEN (kcuda=0, DSR 32 == baseline 32); all four hook checkers OK in the reproduced sandbox; test_async_llm 8/8-347, test_engine_core 6/6-44, test_llm_engine 11/11-204 (CPU still resolves async-ON); clean -Werror CPU build; full 11-gate record battery green. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude (Opus 5) via Claude Code
localai-bot
left a comment
There was a problem hiding this comment.
Thanks for the follow-up to #140 — the dual-GPU FP8 resident direction is right, and the peer-mix idea is one we want. Two things to sort before this can land, one mechanical and one structural. They have the same fix.
1. The CPU, CUDA and Vulkan builds do not link
build-test-cpu and build-test-vulkan are failing for a real reason, not a flake:
gemma4_moe.cpp: undefined reference to `vllm::PinGemma4Fp8ExpertHostCache(...)`
gemma4_moe.cpp: undefined reference to `vllm::PeerCopyGemma4ExpertSlice(...)`
gemma4_moe.cpp: undefined reference to `vt::rocm::MatmulBTFp8ChannelRocm(...)`
gemma4_moe.cpp: undefined reference to `vt::rocm::ExpertGeGLUBf16TopKM1Rocm(...)`
gemma4.cpp: undefined reference to `vt::rocm::RmsNormPlusAddRocm(...)`
gemma4.cpp: undefined reference to `vt::rocm::DualRmsNormPlusResRocm(...)`
gemma4.cpp: undefined reference to `vt::rocm::GeluMulSeparateRocm(...)`
These are declared in include/vllm/model_executor/models/gemma4_moe.h but defined only in src/vt/rocm/rocm_gemma4_experts.hip, and that file is added to the build only under the ROCm branch (CMakeLists.txt:1148). Every non-ROCm configuration compiles the calls and then has nothing to link them against.
2. Backend names in the device-agnostic layer
The deeper reason it breaks is that src/vllm/model_executor/models/ is the shared, device-agnostic layer, and this PR puts 11 vt::rocm:: call sites into it (4 in gemma4.cpp, 7 in gemma4_moe.cpp). Today main has exactly zero — the only vt::rocm:: reference outside src/vt/ is in src/vllm/platforms/rocm.cpp, which is the right place for it. This is what check-device-leakage guards, and it is a rule we hold ourselves to: I had to fix one of my own violations in #159 this week for the same reason.
(Do rerun CI after rebasing, by the way — the device-leakage failure you are seeing right now is partly our bug, not yours. It reports kcuda=1 from a DeviceType::kCUDA I left in runner.cpp. That is fixed on main in #159.)
The fix for both
Route these through the ops/backend seam rather than calling the ROCm entry points by name:
- Kernels (
RmsNormPlusAddRocm,DualRmsNormPlusResRocm,GeluMulSeparateRocm,MatmulBTAlphaBetaRocm,MatmulBTFp8ChannelRocm,ExpertGeGLUBf16TopKM1Rocm): register them asOpIdimplementations forDeviceType::kROCmand call the portablevt::entry point. The model file then reads the same on every backend, and CPU keeps its reference implementation. - Policy/capability (
PeerCopyGemma4ExpertSlice,PinGemma4Fp8ExpertHostCache): these are "can this backend do peer copy / pinned expert staging", so they belong onvt::Backendas virtuals with a conservative base implementation, in the shape ofSupportsAuxStream(include/vt/backend.h:122) andSupportsAsyncSampledTokenReadback(:141). Non-ROCm backends inherit the base and the peer path simply does not engage.
That keeps your ROCm fast path exactly as fast, and makes every other build link again.
Rebase onto current main while you are in there — the merge base here is c05cee1d, several commits back — and agent-record / documentation-checkpoint should be satisfied per AGENTS.md (a record under .agents/ bound to this row, plus the docs/ checkpoint files).
Happy to help with the ops-registration wiring if it would be useful — say the word.
|
Thanks for the clear CHANGES_REQUESTED review — fully agree on both points. Ack
Process / PR hygiene
Rebasing/reworking the ROCm registration path next. Happy to take a pointer if there's a preferred OpId naming pattern for the fused RmsNorm/GeGLU helpers beyond matching existing MatmulBT-style registrations. FOLLOWING_AGENTS_PROTOCOL |
Address mudler#154 review: device-agnostic model layer no longer names ROCm. Fused helpers dispatch ROCm kernels under VLLM_CPP_HIP; non-HIP peer/pin stubs. check-device-leakage holds baseline. Docs STATUS/BENCHMARKS/FEATURES/USAGE. Drop PR_DRAFT_GEMMA4_RESIDENT_SPEED.md from tree (pr-size unclassified path). FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Address mudler#154 review: device-agnostic model layer no longer names ROCm. Fused helpers in include/vt/fused_ops.h (ROCm under VLLM_CPP_HIP; non-HIP stubs). Peer/pin stubs off-HIP. check-device-leakage holds baseline. Drop PR_DRAFT_GEMMA4_RESIDENT_SPEED.md (pr-size). Docs STATUS/BENCHMARKS/FEATURES/USAGE/README. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
e6a314b to
324c405
Compare
Progress per layer, 3GiB headroom on compute GPU, fill GPU0 before GPU1 for same-device fast path. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Resident path dequanted every expert into permanent host caches (~1.5GiB/layer), OOM on ~30G RAM. Stream per-expert ephemeral dequant + H2D; clear caches. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Stream per-expert dequant (no permanent host BF16 cache). 10GiB headroom on compute GPU. Off-device resident layers use ephemeral H2D instead of double-alloc. Lab: 30 layers 42.5GiB resident, Paris stop EXIT=0 peak VRAM ~28.6GB. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Enable HIP peer access; GPU1 expert slices peer/stage to compute scratch instead of FP8 host dequant. Lab: full 30L resident Paris EXIT=0, wall ~198s (was ~294s). FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Default on-GPU weighted accumulate; VT_GEMMA4_HOST_AXPY=1 falls back. Lab: full resident Paris EXIT=0 ~191s wall. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Allocate peer/weight/yscale scratch once per token instead of per expert. Reduces HIP alloc churn during decode. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Stream path no longer reserves ~12MiB expert peer buffers per layer step. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
After one model load, run N blocking completions with per-run secs and tok_s. Lab (Firworks FP8 stream, gfx1201): run1 ~0.14 tok/s (first-use expert dequant), run2 ~24 tok/s warm. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Row-parallel DequantFp8ChannelToBf16; parallel top-k expert cache fill when cold; nesting guard avoids thread storms. Lab: cold ~0.3 tok/s (was ~0.14), warm ~4 tok/s, Paris OK. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
hipHostRegister after first dequant for faster H2D fallback; log device expert OOM fallbacks. Lab: warm still ~3.8-4 tok/s with successful device expert path (Paris OK). FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Shared gate/up/gu/act/weight buffers per token; T=1 pack is two bulk copies. Lab: warm still ~3.8 tok/s (GEMM-bound); Paris OK. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
hipblasGemmStridedBatchedEx + GemmBatchedEx pointer path for top-k expert fuse. Gemma MoE uses VT_GEMMA4_BATCH_EXPERTS=1 to enable; default remains serial ExpertGeGLU (batched path currently slower ~0.7 vs ~3 tok/s warm on R9700 due to Gelu pack). FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Single GeluAndMul over [G,2I] + pointer-batch down AB. Default remains serial (VT_GEMMA4_BATCH_EXPERTS=1 to try). Lab: fused-batch warm ~0.5 tok/s vs serial ~3. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
ExpertScratch, peer/tmp buffers, y/ysum live for the layer call so multi-token prefill and decode avoid re-alloc. Lab warm ~4.5 tok/s (was ~3.9). FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Ungrouped softmax+greedy top-k HIP kernel; Gemma MoE only D2Hs [T,K] weights/indices (not full [T,E] logits). Lab: Paris OK, warm ~4.4 tok/s. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Lab stream FP8 --repeat2: warm mean ~7.6 tok/s (paris 6.8, arith 11), 3/3 correct. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
- GeluMulSeparateRocm: no gate|up pack before GeGLU (ExpertGeGLU hot path) - thread_local MoE scratch reused across 30 layers (pool thrash fix) Lab stream FP8 --repeat2: 3/3 OK, warm mean ~18.9 tok/s (paris 26.7). FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
MatmulBTAlphaBetaRocm: ysum = w*(act@W^T) + beta*ysum — drops MulScalar+Add on device path. Lab: 3/3 OK, warm mean ~18.7 tok/s (paris 24.8). FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Lab: 3/3 OK, warm mean ~19.7 tok/s (paris 27.0). FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Reuse dhn/h1/h2/moe/n*/PLE temps across 30 layers; residual stays on hidden via D2D publish. Lab: 3/3 OK, warm mean ~18.8 tok/s (paris 25.2). FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
- RequestLogger: Received/Finished/api ingress/errors (enable-log-requests) - --enable-log-outputs --max-log-len --enable-metrics - Wire PrometheusStatLogger to GET /metrics - Keep chat-dbg stages under --verbose - Default log-requests + metrics ON for solid agent harness FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Default VT_SERVER_MAX_PROMPT_CHARS=48000, VT_SERVER_MAX_NEW_TOKENS=4096. Prevents async prefill wedge on 140k-char full-agent system dumps. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Pairs with start-gemma4-fp8-8010.sh max_model_len=262144 + prefix cache. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Enabled via VT_SERVER_PREFILL_PROGRESS=1 or VT_SERVER_VERBOSE=1; ~2Hz rate limit. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Root cause of chat hang under concurrent Hermes: LLMEngine::generate looped while has_unfinished_requests() globally, so one huge async job blocked all blocking generate forever. Wait only until this request_id finishes. Chat serving (Gemma4/Hermes, vLLM recipe parity): - apply_chat_template / MakeChatTemplatePromptFn take enable_thinking - --enable-thinking / --no-enable-thinking (default off for agents) - empty thought block when thinking off (HF jinja) Lab: chat count max=48 finishes; warm chat ~32 tok/s after expert cache. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Address mudler#154 review: device-agnostic model layer no longer names ROCm. Fused helpers in include/vt/fused_ops.h (ROCm under VLLM_CPP_HIP; non-HIP stubs). Peer/pin stubs off-HIP. check-device-leakage holds baseline. Drop PR_DRAFT_GEMMA4_RESIDENT_SPEED.md (pr-size). Docs STATUS/BENCHMARKS/FEATURES/USAGE/README. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
324c405 to
a6e1ecf
Compare
Address mudler#154 review: device-agnostic model layer no longer names ROCm. Fused helpers in include/vt/fused_ops.h (ROCm under VLLM_CPP_HIP; non-HIP stubs). Peer/pin stubs off-HIP. check-device-leakage holds baseline. Drop PR_DRAFT_GEMMA4_RESIDENT_SPEED.md (pr-size). Docs STATUS/BENCHMARKS/FEATURES/USAGE/README. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Hermes:grok-4.5 [Hermes]
a6e1ecf to
e086435
Compare
Progress update (post-review)Addressed the structural review items we could land without a full PR split:
Still red (honest)
Split already doneGeneric serve hang fix is #163 (not blocked on this MoE campaign). Next for #154 when size policy is the priority: squash or carve into ≤900-line PRs (ops seam first, then expert path, then kernels). FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true |
localai-bot
left a comment
There was a problem hiding this comment.
Re-reviewed at e086435e — the defect I flagged is fixed, and fixed the right way.
Verified rather than taken on trust:
git grep -c "vt::rocm::" -- src/vllm/model_executor/on your head returns zero (was 11 acrossgemma4.cppandgemma4_moe.cpp).- The new
include/vt/fused_ops.h+src/vt/fused_ops.cppseam carries exactly the right contract in its own header comment: "Model code MUST call these — nevervt::rocm::*— so CPU/CUDA/Vulkan link." Non-HIP stubs sit under the#endif // VLLM_CPP_HIPguard. - CI on that SHA:
build-test-cpu✅,build-test-vulkan✅,device-leakage✅, both sanitizers ✅. The three link failures I quoted are gone.
That is the ops/backend-seam route rather than a #ifdef around the call sites, which is what keeps the device-agnostic layer honest. Nice.
Remaining red is process only — pr-size, agent-record, documentation-checkpoint. Per operator instruction I am merging and carrying those myself rather than bouncing the PR back to you again.
Thanks for the ROCm campaign generally — the gfx1201 evidence in #41 is the strongest data point we have on the discrete lane.
check-public-doc-tables and check-env-doc are red on pristine origin/main and block every push through the pre-push hook. Neither came from a feature branch; both arrived with #154. This fixes the two conditions I can judge correctly and deliberately leaves the rest rather than guessing at another row's intent. FIXED. docs/STATUS.md carried 12 h2 sections against an 11 ratchet. The 12th was a DATED per-change narrative -- exactly what the checker says to collapse -- so its binding result becomes one line under the existing "Backend detail" section, which already owns per-backend state. Nothing is lost: the full narrative stays in the append-only record and the PR. Sections 12 -> 11, STATUS 277416 -> 277217 chars. VT_SERVER_VERBOSE is documented rather than allowlisted because it is user-facing: examples/server sets it from its own verbosity flag, and it is the umbrella switch VT_SERVER_PREFILL_PROGRESS falls back to when unset. NOT FIXED, and named so the next session does not rediscover them: * STATUS is still 277217 against a 276960 ratchet -- main was 456 over before this change and is 257 over after it. Closing that means cutting content I do not own, and the ratchet may only shrink, so it needs the owning row. * VT_GEMMA4_BATCH_EXPERTS, VT_GEMMA4_CUSTOM_EXPERT and VT_GEMMA4_EXPERT_VRAM_MB are undocumented. Each is either a user-facing knob for docs/ENVIRONMENT.md or a kernel-internal switch for the allowlist, and that is a statement about #154's intent. Guessing it would put a wrong claim on a public page, which is worse than leaving the gate red. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude-Code:claude-opus-5 [Claude Code]
docs/BENCHMARKS.md is a KEYED TABLE: a checkpoint updates a ROW, it does not append a dated H2 section. #154 appended one, which left main red on check-public-doc-tables for two reasons at once -- the non-canonical section, and the em-dash it carried (house style forbids em-dashes on the public pages). scripts/roll-benchmark-record.py --apply moves it verbatim into .agents/benchmark-record.md, which is where per-change narrative belongs, and both errors go with it. Nothing is lost and nothing is rewritten: the section moves byte-for-byte into the append-only record. This is main's red, not this branch's, and it was blocking the agent-record CI job on every open PR rather than just on the change that introduced it. RESIDUE, NAMED RATHER THAN PAPERED OVER. check-public-doc-tables still fails on one item: docs/STATUS.md is 277213 chars against a 276960 ratchet. This branch SHRINKS that page by 4 chars; main is already 257 over. Clearing it means collapsing superseded narrative, and the only block big enough to matter is a single 33,211-char table cell (the Laguna-S-2.1 MoE row) that is itself well past the 220-char cell rule. Collapsing a cell that size is a deliberate, separately-reviewable change with its own owner, not something to bury in a Vulkan performance PR, so it is left open and stated here instead. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude-Code:claude-opus-5 [Claude Code]
perf(vulkan): fused attn preamble native -- 16 host round trips/token gone kAttnQkNormRopeGate gets a native Vulkan kernel, removing the last reference-tier op that fires during 27B decode. Module count 24 -> 25; the Vulkan decline list is now exactly kRopeCosSinCache and kCausalConv1dFwd, both host-side by design. WHY THIS OP. After the GDN rows there was no kernel-speed lever left: both decode GEMMs sit near the GB10 bandwidth roof. What remained was the reference tier's ARCHITECTURAL cost -- src/vt/op_provider.cpp drains the recorded command batch (submit plus blocking fence) before it can hand a host kernel device memory, so a reference-tier op costs a full GPU round trip however little arithmetic it does. Of the three ops still declining, kCausalConv1dFwd is prefill-only and kRopeCosSinCache is deliberately on the host (its angle table is built in double precision, vLLM's own split), which left this one. The 27B has 64 layers of which 48 are linear-attention, so it fires 16x per token. PORTED FROM. Per-element math 1:1 from src/vt/cpu/cpu_ops.cpp:956-1010, our transcription of vLLM's QKNormRoPEFusionPass -> _C.fused_qk_norm_rope. Dispatch shape from src/vt/cuda/cuda_ops.cu:1316-1394: one workgroup per (token, head) over Hq + Hkv slots, CUDA's dim3(t, hq+hkv) flattened, the two paired normed elements recomputed per output rather than staged in shared memory -- CUDA's choice, kept so both devices compute the same expression. It is a fusion of two kernels that already work: the reduction is vt_rms_norm_gated.comp's, the cos/sin indexing is vt_rope_from_cache.comp's minus the positions indirection. MEASURED, GB10, 1 prompt, 32-in, c1, page cache dropped before every run. Two-length flush diff (output-len 4 vs 12), both arms measured: per decoded token main this row reference-tier drains 16 0 command-buffer flushes 18 4 GPU dispatches 884 900 GPU-active 236.3 ms 234.6 ms wall (median TPOT) 253.5 ms 246.2 ms The GPU does the same work to within noise and takes four MORE dispatches; what changes is that 16 submit-plus-blocking-fence round trips per token stop happening. Decode over 6 order-alternated AB/BA pairs: this row wins 5 of 6, median TPOT 249.74 -> 242.34 ms = 4.00 -> 4.13 tok/s. Discarding the two pairs carrying an outlier leg leaves 4 of 4 and the same ~3%. llama.cpp Vulkan on this box is 4.35. Attributing the ~3% to the removed round trips is INFERRED: host time is wall minus GPU-active, a derived quantity, not directly instrumented. Gates re-run independently rather than taken on report: on llvmpipe from a clean build, test_vulkan_backend 29/29 (1786 assertions), test_opt_paged_engine on Vulkan 6/6 token-exact (96/96 tokens) with 0 declines, test_backend_cross_device 11/11, and gen-vulkan-spirv.py --check reproduces the committed SPIR-V byte-for-byte under the pinned glslang 16.5.0. CI's own runners pass build-test-vulkan, build-test-cpu and both sanitizers. The f32 arm of the new kernel is NMSE-tier, not bit-exact -- the workgroup tree reduction reorders the mean square, the same trade vt_rms_norm and vt_rms_norm_gated already make; the bf16-q/k arm the model actually uses IS bit-exact. ALSO IN THIS MERGE, and deliberately not silent: docs/BENCHMARKS.md is a keyed table, and #154 had appended a dated H2 section to it. That section plus the em-dash it carried was failing check-public-doc-tables on MAIN, and therefore the agent-record job on every open PR rather than only on the change that caused it. roll-benchmark-record.py --apply moves it verbatim into the append-only record. TWO GATES REMAIN RED AND ARE NAMED RATHER THAN PAPERED OVER. check-pr-size reports the product class at 1306 lines against a 900 budget; 723 of those are the regenerated src/vt/vulkan/vulkan_spirv.cpp, a machine-generated hex blob whose freshness the --check gate proves, so hand-written lines are ~583, inside budget. The checker classifying generated SPIR-V as product is a real gap, left open. check-public-doc-tables still fails on docs/STATUS.md at 277213 chars against a 276960 shrink-only ratchet: this branch SHRINKS that page by 4 chars and main is already 257 over, and the only block big enough to clear it is a single 33,211-char table cell that is itself far past the 220-char cell rule. Collapsing a cell that size is a separately-reviewable change with its own owner. README.md:310 still reads "24 native ops"; check-doc-checkpoint rejects a README change without an accompanying landing-page source, so it needs to ride a change that qualifies. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude-Code:claude-opus-5 [Claude Code]
…lags main has been red on every run overnight. Four of the failures are TREE-scoped, which is the part that matters: unlike the diff-scoped range gates, a tree-scoped failure reds every run AND every open PR until it is fixed, so it blocked unrelated contributor work. 1) check-public-doc-tables: docs/STATUS.md was 277214 chars against a 276960 shrink-only ratchet. The DeepSeek last-mile Bricks 3 and 4 were a run-by-run log of a NOT-pushed branch, 3871 chars for two negative results, on a page whose contract is one binding line per capability. Collapsed to 1009 chars that KEEP both negatives and the finding they establish: dp4a near-flat +0.5%, aligned repack a MEASURED NEGATIVE at +4.7 GiB, coalesced-load REFUTED, the Q8_0 matvec LATENCY/OCCUPANCY- bound at ~61% of ds4. The nsys tables were already in .agents/benchmark-record.md. Ratchet lowered 276960 -> 274760. The 2026-08-08 ratchet left ~15 chars of headroom. That was mine and it was too tight: three Vulkan landings that night each added a status line and pushed the page over within hours. The new headroom is 408, so the ratchet measures bloat rather than merge timing. 2) check-env-doc: 12 production env vars were undocumented -- seven VT_GEMMA4_*, three VT_ROCM_*, and the two VT_SERVER_MAX_* request limits. Documented in docs/ENVIRONMENT.md next to their siblings rather than allowlisted, because the existing VT_GEMMA4_RESIDENT_* and VT_ROCM_ATTN_CPU_REF knobs are documented there. Defaults were read from the source, not guessed. 3) check-state-order: the 2026-08-09T13:00 anchor sat before its heading instead of on the line after it, so that entry counted as unanchored. 4) check-now-current: NOW.md was 6080 chars against a 6000 budget. Dropped the two rows carrying no open work -- MXFP4 parity, marked TERMINAL, and the supported-models list, marked LANDED with a "-" residual. Both keep their evidence in docs/FEATURES.md and the matrices. Also the docs #154 owed, which is what reddened documentation-checkpoint when it landed. It shipped eight real user-facing flags with no USAGE entry: --repeat on vllm-cli, and --enable/--disable-log-requests, --enable-log-outputs, --max-log-len, --enable/--disable-metrics, --enable/--no-enable-thinking and --verbose on the server. NOT documented: --default-chat-template-kwargs, which appears only in a source comment naming vLLM's spelling and is not a parsed flag. README: NOT changed here, and that is the one thing this PR could not fix. The count says 24 native ops in two places where STATUS and NOW both say 25, which .agents/state.md had already flagged. POL-DOC-README only accepts a README edit alongside one of six landing-source files (.agents/mission.md, CMakeLists.txt, benchmarks/demo/*.json, examples/{cli,server}/main.cpp), and a stale-number correction touches none of them and should not have to. Left for the developer to rule on, as state.md asked: either waive it or let the gate accept a correction whose source already sits in docs/BENCHMARKS.md. 5) tests/scripts/test_check_public_doc_tables.py had its `if __name__ == "__main__": unittest.main()` block at line 362, BEFORE the StatusRatchet class at line 392. CI runs this file as a script, so unittest.main() executed before that class was defined and all 8 of its tests -- including test_growth_past_the_char_ratchet_is_rejected and test_the_live_page_is_inside_its_ratchet -- never ran. The ratchet's own mutation suite was inert. Moving the block to the end takes the file from 41 to 49 collected tests; all 49 pass. Found because check-pr-size requires mutation evidence for a checker change and the new test I added did not fail when it should have. It was not failing because it was not running. The new test pins what this change actually alters: ratchet headroom is bounded, so a red char ratchet cannot be cleared by inflating the number instead of shrinking the page. Two-way mutation proof, bytecode disabled: chars=279000 -> FAILED "4648 not less than or equal to 2000"; chars=274000 -> FAILED "-352 not greater than or equal to 0"; baseline 274760 -> OK. Not fixed here, because it is not a bug: POL-PR-REQUIRED keeps firing on content commits from --merge landings and direct pushes to main. That gate is working. The permanent fix is the repository setting (allow squash only), which needs admin. Gates: all 20 tree-scoped checkers green, plus test_agent_record, test_doc_checkpoint, test_check_public_doc_tables, test_check_readme_structure, test_check_env_doc, test_check_state_order, test_check_now_current, test_agent_role, test_agent_gates, test_check_gate_commands and test_audit_live_rows -- on a worktree pinned at f921062. No code touched. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude-Code:claude-opus-5 [Claude Code]
…lags main has been red on every run overnight. Four of the failures are TREE-scoped, which is the part that matters: unlike the diff-scoped range gates, a tree-scoped failure reds every run AND every open PR until it is fixed, so it blocked unrelated contributor work. 1) check-public-doc-tables: docs/STATUS.md was 277214 chars against a 276960 shrink-only ratchet. The DeepSeek last-mile Bricks 3 and 4 were a run-by-run log of a NOT-pushed branch, 3871 chars for two negative results, on a page whose contract is one binding line per capability. Collapsed to 1009 chars that KEEP both negatives and the finding they establish: dp4a near-flat +0.5%, aligned repack a MEASURED NEGATIVE at +4.7 GiB, coalesced-load REFUTED, the Q8_0 matvec LATENCY/OCCUPANCY- bound at ~61% of ds4. The nsys tables were already in .agents/benchmark-record.md. Ratchet lowered 276960 -> 274760. The 2026-08-08 ratchet left ~15 chars of headroom. That was mine and it was too tight: three Vulkan landings that night each added a status line and pushed the page over within hours. The new headroom is 408, so the ratchet measures bloat rather than merge timing. 2) check-env-doc: 12 production env vars were undocumented -- seven VT_GEMMA4_*, three VT_ROCM_*, and the two VT_SERVER_MAX_* request limits. Documented in docs/ENVIRONMENT.md next to their siblings rather than allowlisted, because the existing VT_GEMMA4_RESIDENT_* and VT_ROCM_ATTN_CPU_REF knobs are documented there. Defaults were read from the source, not guessed. 3) check-state-order: the 2026-08-09T13:00 anchor sat before its heading instead of on the line after it, so that entry counted as unanchored. 4) check-now-current: NOW.md was 6080 chars against a 6000 budget. Dropped the two rows carrying no open work -- MXFP4 parity, marked TERMINAL, and the supported-models list, marked LANDED with a "-" residual. Both keep their evidence in docs/FEATURES.md and the matrices. Also the docs #154 owed, which is what reddened documentation-checkpoint when it landed. It shipped eight real user-facing flags with no USAGE entry: --repeat on vllm-cli, and --enable/--disable-log-requests, --enable-log-outputs, --max-log-len, --enable/--disable-metrics, --enable/--no-enable-thinking and --verbose on the server. NOT documented: --default-chat-template-kwargs, which appears only in a source comment naming vLLM's spelling and is not a parsed flag. README: NOT changed here, and that is the one thing this PR could not fix. The count says 24 native ops in two places where STATUS and NOW both say 25, which .agents/state.md had already flagged. POL-DOC-README only accepts a README edit alongside one of six landing-source files (.agents/mission.md, CMakeLists.txt, benchmarks/demo/*.json, examples/{cli,server}/main.cpp), and a stale-number correction touches none of them and should not have to. Left for the developer to rule on, as state.md asked: either waive it or let the gate accept a correction whose source already sits in docs/BENCHMARKS.md. 5) tests/scripts/test_check_public_doc_tables.py had its `if __name__ == "__main__": unittest.main()` block at line 362, BEFORE the StatusRatchet class at line 392. CI runs this file as a script, so unittest.main() executed before that class was defined and all 8 of its tests -- including test_growth_past_the_char_ratchet_is_rejected and test_the_live_page_is_inside_its_ratchet -- never ran. The ratchet's own mutation suite was inert. Moving the block to the end takes the file from 41 to 49 collected tests; all 49 pass. Found because check-pr-size requires mutation evidence for a checker change and the new test I added did not fail when it should have. It was not failing because it was not running. The new test pins what this change actually alters: ratchet headroom is bounded, so a red char ratchet cannot be cleared by inflating the number instead of shrinking the page. Two-way mutation proof, bytecode disabled: chars=279000 -> FAILED "4648 not less than or equal to 2000"; chars=274000 -> FAILED "-352 not greater than or equal to 0"; baseline 274760 -> OK. 6) check-pr-size classified neither CLAUDE.md (added by the b8d996e direct push) nor MANIFESTO.md (a9a8581), so classify_path RAISED on them and any PR touching either failed the gate outright. Both are top-level reader-facing prose: same class as README. Mutation: dropping CLAUDE.md back out gives "Lists differ: ['CLAUDE.md'] != []". Not fixed here, because it is not a bug: POL-PR-REQUIRED keeps firing on content commits from --merge landings and direct pushes to main. That gate is working. The permanent fix is the repository setting (allow squash only), which needs admin. Gates: all 20 tree-scoped checkers green, plus test_agent_record, test_doc_checkpoint, test_check_public_doc_tables, test_check_readme_structure, test_check_env_doc, test_check_state_order, test_check_now_current, test_agent_role, test_agent_gates, test_check_gate_commands and test_audit_live_rows -- on a worktree pinned at f921062. No code touched. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude-Code:claude-opus-5 [Claude Code]
…#154 docs #190 landed while this branch was open and fixed more of main's red than this did, and did it better: the STATUS.md ratchet cleared by MOVING a 33,211-char cell verbatim per POL-EVIDENCE-PRESERVE rather than condensing it, CLAUDE.md/MANIFESTO.md classified (CLAUDE.md as procedure -- it is a symlink to AGENTS.md -- which is more accurate than the public_document I had used), a generated path class for SPIR-V, and a real benchmarks/demo/vulkan_27b_llamacpp.json landing source so the false README could finally be corrected. Everything this branch had that #190 also covers is dropped rather than re-litigated. What was still red on main at c1716fd, both TREE-scoped, so both were failing every run and every open PR: 1. check-agent-record: .agents/parity-ledger.md carried `../examples/server/main.cpp#L90`, and #189 turned that file into a 23-line thin ABI client, so the anchor pointed past EOF. The referenced "server flag" is the CUDA-graph-replay profiler trigger, which now lives at src/vllm/entrypoints/openai/server_main.cpp:692 (`args.cuda_profile_graph_replays > 0`). Repointed there. Only the pointer moves; the ledger entry is untouched. 2. check-env-doc: VLLM_CPP_HTTP_FIXED_POOL, added by #189, was undocumented. Documented in docs/ENVIRONMENT.md rather than allowlisted, because it changes runtime behavior (`=0` reverts the HTTP worker pool to legacy dynamic) and the comparable A/B switches VT_ROCM_ATTN_CPU_REF and VT_DEBUG_SAMPLED are documented there. 3. tests/scripts/test_check_public_doc_tables.py still has its `if __name__ == "__main__": unittest.main()` at line 362, BEFORE the StatusRatchet class at 392. CI runs this file as a script, so unittest.main() executes before that class is defined and all 8 of its tests -- including test_growth_past_the_char_ratchet_is_rejected and test_the_live_page_is_inside_its_ratchet -- never run. The ratchet's own mutation suite is inert, which is notable given #190 just re-pinned that ratchet. Moving the block to the end takes the file from 41 to 49 collected tests; all 49 pass against #190's 244486. 4. docs/USAGE.md: #154 shipped eight user-facing flags with no USAGE entry and #189 did not add them either. Re-verified every one still parses after the thin-client refactor: --repeat in examples/cli/main.cpp, and --enable/--disable-log-requests, --enable-log-outputs, --max-log-len, --enable/--disable-metrics, --enable/--no-enable-thinking and --verbose in src/vllm/entrypoints/openai/server_main.cpp. NOT documented: --default-chat-template-kwargs, which appears only in a source comment naming vLLM's spelling and is not a parsed flag. 5. docs/BENCHMARKS.md still told the reader the Vulkan-vs-llama.cpp arm was "Not yet runnable (no model runs on Vulkan)". It runs; that row now names the harness. 6. docs/STATUS.md said the multimodal server seam "is wired into `examples/server/main.cpp`", which #189 made false -- the `preprocessor_config.json` guard now lives in src/vllm/entrypoints/openai/server_main.cpp (verified by grep, not assumed). Corrected. The longer path costs 19 chars against 3 of ratchet headroom, so it is paid for inside the same cell by tightening two clauses that restate what the sentence already says. Net -5. Note for whoever touches docs/STATUS.md next: #190 tightened the char ratchet to the measured byte, so slack is 3. Any status line added without paying for it reds a tree-scoped gate again. That is a deliberate choice of #190's, left alone here. Gates: 14 tree-scoped checkers green, plus test_check_public_doc_tables (49), test_agent_record, test_doc_checkpoint, test_check_state_order, test_check_now_current, test_check_env_doc and test_check_pr_size (21), on a worktree pinned at c1716fd. No product code touched. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude-Code:claude-opus-5 [Claude Code]
…#154 docs #190 landed while this branch was open and fixed more of main's red than this did, and did it better: the STATUS.md ratchet cleared by MOVING a 33,211-char cell verbatim per POL-EVIDENCE-PRESERVE rather than condensing it, CLAUDE.md/MANIFESTO.md classified (CLAUDE.md as procedure -- it is a symlink to AGENTS.md -- which is more accurate than the public_document I had used), a generated path class for SPIR-V, and a real benchmarks/demo/vulkan_27b_llamacpp.json landing source so the false README could finally be corrected. Everything this branch had that #190 also covers is dropped rather than re-litigated. What was still red on main at c1716fd, both TREE-scoped, so both were failing every run and every open PR: 1. check-agent-record: .agents/parity-ledger.md carried `../examples/server/main.cpp#L90`, and #189 turned that file into a 23-line thin ABI client, so the anchor pointed past EOF. The referenced "server flag" is the CUDA-graph-replay profiler trigger, which now lives at src/vllm/entrypoints/openai/server_main.cpp:692 (`args.cuda_profile_graph_replays > 0`). Repointed there. Only the pointer moves; the ledger entry is untouched. 2. check-env-doc: VLLM_CPP_HTTP_FIXED_POOL, added by #189, was undocumented. Documented in docs/ENVIRONMENT.md rather than allowlisted, because it changes runtime behavior (`=0` reverts the HTTP worker pool to legacy dynamic) and the comparable A/B switches VT_ROCM_ATTN_CPU_REF and VT_DEBUG_SAMPLED are documented there. 3. tests/scripts/test_check_public_doc_tables.py still has its `if __name__ == "__main__": unittest.main()` at line 362, BEFORE the StatusRatchet class at 392. CI runs this file as a script, so unittest.main() executes before that class is defined and all 8 of its tests -- including test_growth_past_the_char_ratchet_is_rejected and test_the_live_page_is_inside_its_ratchet -- never run. The ratchet's own mutation suite is inert, which is notable given #190 just re-pinned that ratchet. Moving the block to the end takes the file from 41 to 49 collected tests; all 49 pass against #190's 244486. 4. docs/USAGE.md: #154 shipped eight user-facing flags with no USAGE entry and #189 did not add them either. Re-verified every one still parses after the thin-client refactor: --repeat in examples/cli/main.cpp, and --enable/--disable-log-requests, --enable-log-outputs, --max-log-len, --enable/--disable-metrics, --enable/--no-enable-thinking and --verbose in src/vllm/entrypoints/openai/server_main.cpp. NOT documented: --default-chat-template-kwargs, which appears only in a source comment naming vLLM's spelling and is not a parsed flag. 5. docs/BENCHMARKS.md still told the reader the Vulkan-vs-llama.cpp arm was "Not yet runnable (no model runs on Vulkan)". It runs; that row now names the harness. 6. docs/STATUS.md said the multimodal server seam "is wired into `examples/server/main.cpp`", which #189 made false -- the `preprocessor_config.json` guard now lives in src/vllm/entrypoints/openai/server_main.cpp (verified by grep, not assumed). Corrected. The longer path costs 19 chars against 3 of ratchet headroom, so it is paid for inside the same cell by tightening two clauses that restate what the sentence already says. Net -5. Note for whoever touches docs/STATUS.md next: #190 tightened the char ratchet to the measured byte, so slack is 3. Any status line added without paying for it reds a tree-scoped gate again. That is a deliberate choice of #190's, left alone here. Gates: 14 tree-scoped checkers green, plus test_check_public_doc_tables (49), test_agent_record, test_doc_checkpoint, test_check_state_order, test_check_now_current, test_check_env_doc and test_check_pr_size (21), on a worktree pinned at c1716fd. No product code touched. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude-Code:claude-opus-5 [Claude Code]
…#154 docs (#188) #190 landed while this branch was open and fixed more of main's red than this did, and did it better: the STATUS.md ratchet cleared by MOVING a 33,211-char cell verbatim per POL-EVIDENCE-PRESERVE rather than condensing it, CLAUDE.md/MANIFESTO.md classified (CLAUDE.md as procedure -- it is a symlink to AGENTS.md -- which is more accurate than the public_document I had used), a generated path class for SPIR-V, and a real benchmarks/demo/vulkan_27b_llamacpp.json landing source so the false README could finally be corrected. Everything this branch had that #190 also covers is dropped rather than re-litigated. What was still red on main at c1716fd, both TREE-scoped, so both were failing every run and every open PR: 1. check-agent-record: .agents/parity-ledger.md carried `../examples/server/main.cpp#L90`, and #189 turned that file into a 23-line thin ABI client, so the anchor pointed past EOF. The referenced "server flag" is the CUDA-graph-replay profiler trigger, which now lives at src/vllm/entrypoints/openai/server_main.cpp:692 (`args.cuda_profile_graph_replays > 0`). Repointed there. Only the pointer moves; the ledger entry is untouched. 2. check-env-doc: VLLM_CPP_HTTP_FIXED_POOL, added by #189, was undocumented. Documented in docs/ENVIRONMENT.md rather than allowlisted, because it changes runtime behavior (`=0` reverts the HTTP worker pool to legacy dynamic) and the comparable A/B switches VT_ROCM_ATTN_CPU_REF and VT_DEBUG_SAMPLED are documented there. 3. tests/scripts/test_check_public_doc_tables.py still has its `if __name__ == "__main__": unittest.main()` at line 362, BEFORE the StatusRatchet class at 392. CI runs this file as a script, so unittest.main() executes before that class is defined and all 8 of its tests -- including test_growth_past_the_char_ratchet_is_rejected and test_the_live_page_is_inside_its_ratchet -- never run. The ratchet's own mutation suite is inert, which is notable given #190 just re-pinned that ratchet. Moving the block to the end takes the file from 41 to 49 collected tests; all 49 pass against #190's 244486. 4. docs/USAGE.md: #154 shipped eight user-facing flags with no USAGE entry and #189 did not add them either. Re-verified every one still parses after the thin-client refactor: --repeat in examples/cli/main.cpp, and --enable/--disable-log-requests, --enable-log-outputs, --max-log-len, --enable/--disable-metrics, --enable/--no-enable-thinking and --verbose in src/vllm/entrypoints/openai/server_main.cpp. NOT documented: --default-chat-template-kwargs, which appears only in a source comment naming vLLM's spelling and is not a parsed flag. 5. docs/BENCHMARKS.md still told the reader the Vulkan-vs-llama.cpp arm was "Not yet runnable (no model runs on Vulkan)". It runs; that row now names the harness. 6. docs/STATUS.md said the multimodal server seam "is wired into `examples/server/main.cpp`", which #189 made false -- the `preprocessor_config.json` guard now lives in src/vllm/entrypoints/openai/server_main.cpp (verified by grep, not assumed). Corrected. The longer path costs 19 chars against 3 of ratchet headroom, so it is paid for inside the same cell by tightening two clauses that restate what the sentence already says. Net -5. Note for whoever touches docs/STATUS.md next: #190 tightened the char ratchet to the measured byte, so slack is 3. Any status line added without paying for it reds a tree-scoped gate again. That is a deliberate choice of #190's, left alone here. Gates: 14 tree-scoped checkers green, plus test_check_public_doc_tables (49), test_agent_record, test_doc_checkpoint, test_check_state_order, test_check_now_current, test_check_env_doc and test_check_pr_size (21), on a worktree pinned at c1716fd. No product code touched. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude-Code:claude-opus-5 [Claude Code] Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
…#154 docs #190 landed while this branch was open and fixed more of main's red than this did, and did it better: the STATUS.md ratchet cleared by MOVING a 33,211-char cell verbatim per POL-EVIDENCE-PRESERVE rather than condensing it, CLAUDE.md/MANIFESTO.md classified (CLAUDE.md as procedure -- it is a symlink to AGENTS.md -- which is more accurate than the public_document I had used), a generated path class for SPIR-V, and a real benchmarks/demo/vulkan_27b_llamacpp.json landing source so the false README could finally be corrected. Everything this branch had that #190 also covers is dropped rather than re-litigated. What was still red on main at c1716fd, both TREE-scoped, so both were failing every run and every open PR: 1. check-agent-record: .agents/parity-ledger.md carried `../examples/server/main.cpp#L90`, and #189 turned that file into a 23-line thin ABI client, so the anchor pointed past EOF. The referenced "server flag" is the CUDA-graph-replay profiler trigger, which now lives at src/vllm/entrypoints/openai/server_main.cpp:692 (`args.cuda_profile_graph_replays > 0`). Repointed there. Only the pointer moves; the ledger entry is untouched. 2. check-env-doc: VLLM_CPP_HTTP_FIXED_POOL, added by #189, was undocumented. Documented in docs/ENVIRONMENT.md rather than allowlisted, because it changes runtime behavior (`=0` reverts the HTTP worker pool to legacy dynamic) and the comparable A/B switches VT_ROCM_ATTN_CPU_REF and VT_DEBUG_SAMPLED are documented there. 3. tests/scripts/test_check_public_doc_tables.py still has its `if __name__ == "__main__": unittest.main()` at line 362, BEFORE the StatusRatchet class at 392. CI runs this file as a script, so unittest.main() executes before that class is defined and all 8 of its tests -- including test_growth_past_the_char_ratchet_is_rejected and test_the_live_page_is_inside_its_ratchet -- never run. The ratchet's own mutation suite is inert, which is notable given #190 just re-pinned that ratchet. Moving the block to the end takes the file from 41 to 49 collected tests; all 49 pass against #190's 244486. 4. docs/USAGE.md: #154 shipped eight user-facing flags with no USAGE entry and #189 did not add them either. Re-verified every one still parses after the thin-client refactor: --repeat in examples/cli/main.cpp, and --enable/--disable-log-requests, --enable-log-outputs, --max-log-len, --enable/--disable-metrics, --enable/--no-enable-thinking and --verbose in src/vllm/entrypoints/openai/server_main.cpp. NOT documented: --default-chat-template-kwargs, which appears only in a source comment naming vLLM's spelling and is not a parsed flag. 5. docs/BENCHMARKS.md still told the reader the Vulkan-vs-llama.cpp arm was "Not yet runnable (no model runs on Vulkan)". It runs; that row now names the harness. 6. docs/STATUS.md said the multimodal server seam "is wired into `examples/server/main.cpp`", which #189 made false -- the `preprocessor_config.json` guard now lives in src/vllm/entrypoints/openai/server_main.cpp (verified by grep, not assumed). Corrected. The longer path costs 19 chars against 3 of ratchet headroom, so it is paid for inside the same cell by tightening two clauses that restate what the sentence already says. Net -5. 7. check-device-leakage went red on main too, a fifth tree-scoped gate: DSR bucket vt_ifdef 37 > baseline 32. Cause is mechanical -- #189 moved the server body from examples/server/main.cpp, which the scanner never looked at, into src/vllm/entrypoints/openai/server_main.cpp, which IS the shared layer, carrying its 5 `#ifdef VT_BENCH_PROFILE_CONTROL` sites with it. The code did not change. Exempted per-site with the checker's own `// DSR-ALLOW(<row-id>):` hatch rather than a per-file ALLOWLIST budget, for two reasons. A budget can be spent on something else -- swap one guard for a real device fork and the count still reads 5 -- while a per-site marker names each one. And the ALLOWLIST route required editing check-device-leakage.py, which check-pr-size then rejected with "has no affected POL rule mapping": that checker is named by no POL rule in policy.csv, so registering it is a policy decision, not a CI repair. The baseline is NOT raised; DSR returns to 32 == 32 and the 5 exemptions print in CI output every run. The real repair is to put the profiler seam behind a Platform capability query like every other device fork, which is a change to #189's TU. 8. check-pr-size could never accept a change to check-device-leakage.py. recognized_evidence derives tests/scripts/test_check_<name>.py, but that suite predates the convention and CI runs it as tests/scripts/test_device_leakage.py, so every change to that checker failed with "requires semantic mutation evidence in <a file that does not exist>" -- which is exactly what happened above. Added to CHECKER_EVIDENCE_OVERRIDES, which exists for this. An audit of all 30 checkers found one more: check-dsv4-gguf-namemap.py has no suite at all and is not in ci.yml, named in a KNOWN_UNTESTED set so the gap is visible in a test rather than invisible in a naming rule. Mutations: removing the override gives "Lists differ: ['scripts/check-device-leakage.py -> ...'] != []"; downgrading one DSR-ALLOW to a plain comment fails both the new per-site test ("line 50 has no DSR-ALLOW on it or directly above it") and the DSR ratchet (33 > 32). Note for whoever touches docs/STATUS.md next: #190 tightened the char ratchet to the measured byte, so slack is 3. Any status line added without paying for it reds a tree-scoped gate again. That is a deliberate choice of #190's, left alone here. Gates: 14 tree-scoped checkers green, plus test_check_public_doc_tables (49), test_agent_record, test_doc_checkpoint, test_check_state_order, test_check_now_current, test_check_env_doc and test_check_pr_size (21), on a worktree pinned at c1716fd. No product code touched. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude-Code:claude-opus-5 [Claude Code]
…#154 docs #190 landed while this branch was open and fixed more of main's red than this did, and did it better: the STATUS.md ratchet cleared by MOVING a 33,211-char cell verbatim per POL-EVIDENCE-PRESERVE rather than condensing it, CLAUDE.md/MANIFESTO.md classified (CLAUDE.md as procedure -- it is a symlink to AGENTS.md -- which is more accurate than the public_document I had used), a generated path class for SPIR-V, and a real benchmarks/demo/vulkan_27b_llamacpp.json landing source so the false README could finally be corrected. Everything this branch had that #190 also covers is dropped rather than re-litigated. What was still red on main at c1716fd, both TREE-scoped, so both were failing every run and every open PR: 1. check-agent-record: .agents/parity-ledger.md carried `../examples/server/main.cpp#L90`, and #189 turned that file into a 23-line thin ABI client, so the anchor pointed past EOF. The referenced "server flag" is the CUDA-graph-replay profiler trigger, which now lives at src/vllm/entrypoints/openai/server_main.cpp:692 (`args.cuda_profile_graph_replays > 0`). Repointed there. Only the pointer moves; the ledger entry is untouched. 2. check-env-doc: VLLM_CPP_HTTP_FIXED_POOL, added by #189, was undocumented. Documented in docs/ENVIRONMENT.md rather than allowlisted, because it changes runtime behavior (`=0` reverts the HTTP worker pool to legacy dynamic) and the comparable A/B switches VT_ROCM_ATTN_CPU_REF and VT_DEBUG_SAMPLED are documented there. 3. tests/scripts/test_check_public_doc_tables.py still has its `if __name__ == "__main__": unittest.main()` at line 362, BEFORE the StatusRatchet class at 392. CI runs this file as a script, so unittest.main() executes before that class is defined and all 8 of its tests -- including test_growth_past_the_char_ratchet_is_rejected and test_the_live_page_is_inside_its_ratchet -- never run. The ratchet's own mutation suite is inert, which is notable given #190 just re-pinned that ratchet. Moving the block to the end takes the file from 41 to 49 collected tests; all 49 pass against #190's 244486. 4. docs/USAGE.md: #154 shipped eight user-facing flags with no USAGE entry and #189 did not add them either. Re-verified every one still parses after the thin-client refactor: --repeat in examples/cli/main.cpp, and --enable/--disable-log-requests, --enable-log-outputs, --max-log-len, --enable/--disable-metrics, --enable/--no-enable-thinking and --verbose in src/vllm/entrypoints/openai/server_main.cpp. NOT documented: --default-chat-template-kwargs, which appears only in a source comment naming vLLM's spelling and is not a parsed flag. 5. docs/BENCHMARKS.md still told the reader the Vulkan-vs-llama.cpp arm was "Not yet runnable (no model runs on Vulkan)". It runs; that row now names the harness. 6. docs/STATUS.md said the multimodal server seam "is wired into `examples/server/main.cpp`", which #189 made false -- the `preprocessor_config.json` guard now lives in src/vllm/entrypoints/openai/server_main.cpp (verified by grep, not assumed). Corrected. The longer path costs 19 chars against 3 of ratchet headroom, so it is paid for inside the same cell by tightening two clauses that restate what the sentence already says. Net -5. 7. check-device-leakage went red on main too, a fifth tree-scoped gate: DSR bucket vt_ifdef 37 > baseline 32. Cause is mechanical -- #189 moved the server body from examples/server/main.cpp, which the scanner never looked at, into src/vllm/entrypoints/openai/server_main.cpp, which IS the shared layer, carrying its 5 `#ifdef VT_BENCH_PROFILE_CONTROL` sites with it. The code did not change. Exempted per-site with the checker's own `// DSR-ALLOW(<row-id>):` hatch rather than a per-file ALLOWLIST budget, for two reasons. A budget can be spent on something else -- swap one guard for a real device fork and the count still reads 5 -- while a per-site marker names each one. And the ALLOWLIST route required editing check-device-leakage.py, which check-pr-size then rejected with "has no affected POL rule mapping": that checker is named by no POL rule in policy.csv, so registering it is a policy decision, not a CI repair. The baseline is NOT raised; DSR returns to 32 == 32 and the 5 exemptions print in CI output every run. The real repair is to put the profiler seam behind a Platform capability query like every other device fork, which is a change to #189's TU. 8. check-pr-size could never accept a change to check-device-leakage.py. recognized_evidence derives tests/scripts/test_check_<name>.py, but that suite predates the convention and CI runs it as tests/scripts/test_device_leakage.py, so every change to that checker failed with "requires semantic mutation evidence in <a file that does not exist>" -- which is exactly what happened above. Added to CHECKER_EVIDENCE_OVERRIDES, which exists for this. An audit of all 30 checkers found one more: check-dsv4-gguf-namemap.py has no suite at all and is not in ci.yml, named in a KNOWN_UNTESTED set so the gap is visible in a test rather than invisible in a naming rule. Mutations: removing the override gives "Lists differ: ['scripts/check-device-leakage.py -> ...'] != []"; downgrading one DSR-ALLOW to a plain comment fails both the new per-site test ("line 50 has no DSR-ALLOW on it or directly above it") and the DSR ratchet (33 > 32). Note for whoever touches docs/STATUS.md next: #190 tightened the char ratchet to the measured byte, so slack is 3. Any status line added without paying for it reds a tree-scoped gate again. That is a deliberate choice of #190's, left alone here. Gates: 14 tree-scoped checkers green, plus test_check_public_doc_tables (49), test_agent_record, test_doc_checkpoint, test_check_state_order, test_check_now_current, test_check_env_doc and test_check_pr_size (21), on a worktree pinned at c1716fd. No product code touched. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude-Code:claude-opus-5 [Claude Code]
…ng (#267) Merges `row/SAMPLE-LOGPROB-TOKEN-IDS` `d6b482b6`, closing #264. Spec `.agents/specs/logprob-token-ids.md`, committed in its own commit before any implementation code. Row `SAMPLE-LOGPROB-TOKEN-IDS` stays `PARTIAL`. WHAT IT DOES. A request may now name an EXPLICIT set of vocab ids and get back exactly those plus the sampled token -- vLLM's generative-scoring path (`generative_scoring/serving.py:247-255` sends `max_tokens=1` + `logprob_token_ids=label_token_ids`), and the efficient answer to "score these five labels" without `logprobs=-1` and a full-vocab sort. Ported 1:1 at pin `555967922`: the field and `MAX_LOGPROB_TOKEN_IDS` (`sampling_params.py:31, 278-283`), the `num_logprobs` PROPERTY (`:724-729`), the two config-free validations (`:773-782,795-801`), req-id-keyed `InputBatch` tracking re-keyed to req-index over the live batch (`gpu_input_batch.py:273,443-444,574,934-951`), `gather_specific_token_logprobs` (`sampler.py:151-225`) with the padded `[n, max+1]` row, sampled token in column 0, `-inf` padding and the sampled token's rank over the FULL vocab, the snapshot condition (`:86`) and the precedence rule that explicit ids WIN over a count (`:133-136`). THE HIDDEN HALF, which is what the review loop surfaced. Three consumers spelled upstream's `num_logprobs` PROPERTY as the raw `logprobs` field -- the scheduler's slice gate (`scheduler.py:1818`), `LogprobsProcessor::FromNewRequest` (whose comment already CLAIMED it was the property) and `RequestState::FromNewRequest`. Identical for every request that exists today, so nothing was visibly broken; without fixing them a scoring request produced sampler output that nothing downstream ever read and the feature was unreachable through the engine. A SECOND RED landed mid-implementation for exactly that reason: with the sampler and `InputBatch` fully wired, `test_llm_engine` was still red because `RequestState::FromNewRequest` gated the whole `LogprobsProcessor` on `sp.logprobs.has_value()`, where upstream constructs it unconditionally (`output_processor.py:223-229`). Our device-resident greedy fast path is OURS, not upstream's, and skips the gather entirely; it is now gated on the same combined predicate, and missing that second edit would have made the feature vanish silently on the async greedy path. Issue #249's defect class is kept out (every requested id bounds-checked into `[0, vocab)` and every map key into `[0, n)`, mirroring what `torch.gather` raises on) while its instance, `GatherLogprobs`' unbounded `k`, is deliberately untouched. The branch's own main-merge found three PRODUCT conflicts with #238 (`logprobs_mode`) and #223 (`prompt_logprobs`), which landed after it was written, and resolved all three by keeping BOTH features -- the mode selects WHICH tensor the snapshot holds, the ids select WHICH entries are read out of it. It added ONE test neither PR's base could have carried, "logprob_token_ids reads the PROCESSED snapshot under a processed mode", RED-first proven by mutation (weakening the request to `num_logprobs.has_value() && processed_mode` leaves 20 of 21 cases green and fails only this one), tree restored byte-for-byte. KEYED RECORDS -- main's version taken WHOLESALE, the branch's scoped edit reapplied by hand, every deleted anchor asserted to occur exactly once, and every non-reconciled path proven byte-identical to the branch's own edit set. This is the THIRD of three merges in one landing, so several records had already moved under PR #324 and PR #282 and none of the branch's own measurements survived. docs/STATUS.md auto-merge accepted only after proving the edit set matches the branch's exactly; the Sampling row and #282's LoRA row are different rows. Merged page RE-MEASURED: 243,119 = 243,128 after #282, less the Sampling row's own net -9. STATUS ratchet `scripts/check-public-doc-tables.py` CONFLICT for the second time in this landing. ALL rationale histories kept, append-only: main's, #282's, and #267's. Pinned to the RE-MEASURED 243119. The branch's 243278 was DISCARDED -- it was measured against `5812b8b6`, before main's own re-pin AND before #282's collapse. Byte-tight (`test_the_rebased_character_ratchet_is_byte_tight` requires cap == len(page) exactly), strictly DOWN 243188 -> 243128 -> 243119 across the landing. ratchet CEILING `tests/scripts/test_check_public_doc_tables.py` CONFLICT. Lowered to 243119 in the SAME change as the ratchet; both branches' stale ceilings (243482, 243479 -- both computed against `5812b8b6`) discarded, and both collapses' rationale carried in the comment. EVERY mutation guard from every side is kept: main's `test_the_ratchet_is_exactly_one_byte_wide` and `test_a_repin_can_only_tighten_the_char_ratchet`, this branch's NEW `test_one_char_of_growth_on_the_LIVE_page_is_rejected`, the byte-tight guard, the no-hidden-headroom guard, the only-ever-moves-down guard, and BOTH paid-for-by-a-real-collapse guards (#223's beam-search collapse and #238's best_of rationale). The branch could not re-spend #223's collapse, so it collapsed FOUR different restatements in the same Sampling cell instead -- each a definition of what an OpenAI field DOES rather than a statement of what we support -- and all three substrings those two guards pin are intact. 59 tests, all pass. docs/BENCHMARKS.md the branch's terse 145-byte `logprob_token_ids` row landed inside the headroom PR #282's merge bought by moving the superseded 2026-08-08 `BENCH-VK-LLAMA` row into `.agents/benchmark-record.md`. No cap was raised and nothing else moved for this merge: 44,709 -> 44,854 of the hard 45,000. .agents/NOW.md CONFLICT. Main + #282 carried the `logprobs_mode` (#238) row and #282's compacted `Surface coverage` cell and its new `LORA-RUNTIME` row; this branch UPDATES the #238 row in place into `SAMPLE-LOGPROB-TOKEN-IDS` rather than adding a second one. All three kept. The net +19 took the page to 6,005 over the hard 6,000 budget, so it was paid for inside the page: the `Work:` line still announced the PREVIOUS landing and now names this one, and the TP spike row's "(unblocks #127/#154/#155)" clause is stale -- #154 and #155 are MERGED and #127 is CLOSED, verified with `gh pr view`. 5,978 chars / 94 lines. .agents/engine-matrix.md the `SAMPLE-LOGPROB-TOKEN-IDS` row auto-merged into main's file and was verified line-for-line against the branch's scoped edit. The row's State does NOT move (`PARTIAL` before and after -- #238 already moved it off `INVENTORIED`), only its Owner gains `CLAIM-SAMPLE-LOGPROB-TOKEN-IDS`, so the lifecycle rollup is deliberately untouched here and stays at the totals PR #282's merge RECOMPUTED. `check-agent-record` confirms ENGINE=147. .agents/coordination.md the branch's prose claim only; the claims TABLE that `check-agent-record` cross-references is byte-for-byte main's plus #282's `CLAIM-LORA-RUNTIME-W2` row. .agents/roadmap_v1.md CONFLICT. #282's `#278` issue row and this branch's `#264` row are distinct keys, unioned. The C7 portfolio row's stale gap list is corrected by the branch's own scoped edit. .agents/porting-inventory.md `logprob_token_ids` leaves the deferred-stub list; the same sentence's stale `logprobs_mode` entry is corrected in passing because this edit rewrites it. docs/USAGE.md the new field with an example; #282's LoRA paragraph sits above it and both survive. .agents/benchmark-record.md untouched by this branch beyond the row PR #282's merge moved in; it is the one genuinely append-only log in this landing. RESIDUALS, which is why the row is `PARTIAL`: the `logprobs_mode` variants (open PR #258 owns them), the OpenAI request field on `/v1/completions`, `/v1/chat/completions` and `/v1/generative_scoring`, and vocab-range validation in `Verify()` -- which has no model config, exactly as for `allowed_token_ids`, and the sampler bounds the ids anyway, so that one is message quality, not safety. GATE, re-run by the operator on the merged tree, CPU Release, foreground, unbounded: cmake --build build-cpu -j 18 834 targets, 0 errors ./build-cpu/tests/test_sampler 21/21 cases, 114 assertions, 0 skipped ./build-cpu/tests/test_input_batch 29/29 cases, 205 assertions, 0 skipped ./build-cpu/tests/test_llm_engine 24/24 cases, 492 assertions, 0 skipped scripts/check-agent-record.py OK, ENGINE=147 MODEL=362 QUANT=82 KERNEL=51 BACKEND=80 scripts/check-public-doc-tables.py OK scripts/check-now-current.py OK scripts/check-fusion-consistency.py OK, 0 drift tests/scripts/test_check_public_doc_tables.py 59/59 `scripts/__pycache__` was cleared before every checker run: the ratchet values in this landing are the same byte length as the ones they replace, so a stale `.pyc` survives mtime/size invalidation and a checker will silently keep reading the old number. `test_llm_engine` reads 492 assertions here against the 493 the branch measured on its own base. Both are 24/24 cases, 0 failed, 0 skipped -- no case is unreached, so this is not a killed run. The delta is main's, not this merge's: `origin/main` changed sixteen files under `src/vllm/v1/engine` and `include/vllm/v1` between `5812b8b6` (the branch's base) and `91763643`, and its only edit to the test file itself was to ADD five lines. The FULL `ctest -j 6` for this landing is reported on the reconciliation merge that follows this one, because `origin/main` advanced to `75a29016` while these three were being gated and the binding gate belongs on the tree that is actually pushed. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: ClaudeCode:claude-opus-5 [ClaudeCode]
…retired under us `origin/main` advanced a SECOND time during this landing, 31 commits, while the reconciliation merge below was being gated. A plain `git push` was refused as non-fast-forward, which is git protecting those merges, so this fetches, re-merges, reconciles every record against the shape main now has, and re-gates. No force, no rebase of a published branch. THE TWO CONSTRAINTS THIS LANDING WAS BUILT AROUND NO LONGER EXIST. #364/#368 (`ENG-RECORD-CONFLICT-SURFACES`) retired them as a defect, and AGENTS.md gained the rule behind it: *no surface that every PR must write*. Concretely: - `STATUS_RATCHET["chars"]` is DELETED from `scripts/check-public-doc-tables.py`. It was a byte count of one file stored in another and allowed only to fall, so every PR owing STATUS.md a line had to evict unrelated prose AND edit the checker. The three QUALITY keys (`h2_sections`, `long_paragraphs`, `oversized_cells`) are kept. - `MAX_CHARS` is DELETED from `scripts/check-now-current.py`. NOW.md was a fixed-size shared buffer at exactly 6000/6000, so adding a row meant evicting someone else's. `MAX_LINES` and `MAX_ENTRY_CHARS` are kept, and they cap the ENTRY rather than the file. So the ratchet arithmetic this landing carried out three times -- 243188 -> 243128 -> 243119 -> 243117, each byte-tight and re-measured -- is now moot, and the honest resolution is to DROP it rather than defend it. Main's version of both checker files and of `tests/scripts/test_check_public_doc_tables.py` is taken BYTE-FOR-BYTE, which also drops PR #267's `test_one_char_of_growth_on_the_LIVE_page_is_rejected`: that guard asserts on the `chars` key, and the key is gone. Its subject was retired, not its argument. `grep 'STATUS_RATCHET\["chars"\]'` over `scripts/` and `tests/scripts/` returns nothing. Equally, the NOW.md compactions this landing made to buy room -- the `SERVE-METRICS` restatement, the MiniMax-H3 and Release cells, the TP row's `(unblocks #127/#154/#155)` clause -- are reverted to main's fuller text: they were payments for a budget that no longer exists, and carrying them would be unrelated churn in a keyed record. Also fixed on main and no longer owed by anyone: the four conflict markers committed to `.agents/specs/sm120-qwen35-conv-channel-tile-2026-08-08.md`, which this landing found on main and did not touch (`76dfe8dc`). KEYED RECORDS, reconciled against the shape main now has. scripts/check-public-doc-tables.py, tests/scripts/test_check_public_doc_tables.py main's version taken BYTE-FOR-BYTE (`git diff` vs `f64f2b71` is empty for both). The landing's ratchet work is dropped entirely, per above. .agents/roadmap_v1.md CONFLICT. Main SORTED the whole issue table by Row then issue number, which moved every line. Main's sorted table taken wholesale and the landing's SIX rows re-inserted in sorted position -- #299, #314, #337, #338 under `ROAD-V1-C1`, #278 under `LORA-RUNTIME`, #264 beside #238 under `SAMPLE-LOGPROB-TOKEN-IDS`. Verified: all 48 of main's rows present, exactly 6 added, and the only other line that differs from main is the `ROAD-V1-C7` portfolio row, which PR #267 deliberately updates. .agents/engine-matrix.md CONFLICT. Main's `Serving, API, CLI, library` row gained a row and a `READY` (27 -> 28). Main's row kept wholesale, the landing's `LoRA and adapters` row (`ANCHOR-BACKFILL` -> `ACTIVE`) reapplied, and the **Total** RECOMPUTED from the ten actual area rows rather than carried from either side: 148|35|16|4|9|27|8|9|39. `check-agent-record.py` confirms ENGINE=148. .agents/coordination.md auto-merged and verified byte-identical to main's edit set. #368 introduced `.agents/claims/` (one file per claim) but deliberately does NOT migrate existing rows -- the checker still reads the legacy table -- so PR #282's `CLAIM-LORA-RUNTIME-W2` row stays where it is and is removed when the claim closes. scripts/check-gate-commands.py CONFLICT: both sides added a comment block at the same anchor. Unioned; `RUNNABLE_BASELINE` now carries `ENG-RECORD-CONFLICT-SURFACES` (main's), `ENG-RELEASE-CONTAINERS` (main's) and `LORA-RUNTIME` (#282's). 117 gated rows, 32 runnable. docs/FEATURES.md CONFLICT. Main's newer prose wins (37 registered architectures, up from 35) and PR #324's merged-GEMM sentence is reapplied into it. First attempt as a separate paragraph tripped the 21-vs-20 prose-paragraph cap and appending it inline tripped the 700-char paragraph cap at 782, so it is paid for INSIDE the paragraph exactly as #324 originally did -- the lead-in collapsed, total 697 of 700. No cap was raised. .agents/NOW.md CONFLICT. Main's three fuller rows taken wholesale (`BACKEND-ROCM` now records the gfx1100 GDN slice and #269 M0-M4). 94 of 100 lines, every entry inside MAX_ENTRY_CHARS. docs/STATUS.md, docs/BENCHMARKS.md untouched by main this time; 243,117 and 44,839. BENCHMARKS is still 161 chars inside its hard 45,000 cap thanks to the row PR #282's merge moved into `.agents/benchmark-record.md`. GATE for the WHOLE landing, re-run by the operator on THIS tree, CPU Release, foreground, no timeout on any test binary. cmake --build build-cpu -j 18 692 targets, 0 errors, 0 warnings Focused, all seven declared gates of the three PRs: test_dense_gate_up_seam_forward 4/4 cases | 1940 assertions | 0 skipped test_linear_method 6/6 | 76 | 0 skipped test_lora_layers 16/16 | 4498 | 0 skipped test_punica_cpu 8/8 | 149 | 0 skipped test_sampler 21/21 | 114 | 0 skipped test_input_batch 29/29 | 205 | 0 skipped test_llm_engine 24/24 | 494 | 0 skipped Checkers, `scripts/__pycache__` cleared before each: check-agent-record.py OK ENGINE=148 MODEL=362 QUANT=82 KERNEL=51 BACKEND=80 check-public-doc-tables.py OK BENCHMARKS 44,839 / 45,000 FEATURES 27,371, longest prose 697 / 700 check-now-current.py OK 94 / 100 lines check-fusion-consistency.py OK glue 14 / 12 routed / 2 allowlisted; merged-gemm 10 / 6 / 6, 0 drift check-gate-commands.py OK 117 gated rows, 32 runnable check-commit-trailers.py OK over the whole landing range Checker unit suites, all seven: test_check_public_doc_tables 52/52, test_check_gate_commands 28/28, test_check_fusion_consistency 20/20, test_agent_record 29/29, test_check_now_current 11/11, test_record_merge_shape 11/11, test_check_pr_size 38/38. Full `ctest --test-dir build-cpu -j 6 --output-on-failure`, 383 tests, 2771.26 s. BOTH numbers reported: 99% tests passed, 2 failed out of 383 -- `test_async_llm` (0.33 s) and `test_openai_conformance` (163.91 s). Both are on the known starvation-prone list, and both pass SERIALLY on an idle box (load 2.52): test_async_llm 15/15 cases | 443 assertions | wall 0.05 s test_openai_conformance 23/23 | 252 | wall 0.45 s `test_openai_api_server`, which needed a serial re-run last time, PASSED under -j 6 here in 128.64 s -- the same binary, the same tree, a quieter box, which is the clearest available evidence that these are scheduling artifacts rather than defects. Neither failing binary loads a file this landing touches. The #274 ASan/UBSan five did not appear: this gate is Release with no sanitizer, and `test_llm_engine`, `test_capi` and `test_llama_embedding_fold` -- three of that five -- all PASS here (2006.17 s, 1618.64 s, 14.20 s). Note on `test_llm_engine`'s assertion count, corrected earlier in this landing: it is run-to-run NONDETERMINISTIC in this binary (493, 493, 492 measured across three consecutive runs of one unmodified build; 494 here). Every run is 24/24 cases, 0 failed, 0 skipped. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: ClaudeCode:claude-opus-5 [ClaudeCode]
… off ROCm Gemma-4 text generation threw on every non-ROCm build: vt::GeluMulSeparate: ROCm-only fast path in this build gemma4.cpp's `ple > 0` block calls this from the SHARED forward, not from a ROCm branch, and vt::GeluMulSeparate had exactly one implementation guarded by `#if defined(VLLM_CPP_HIP)` with a throw as the fallback. Gemma-4 E4B has per-layer embeddings, so the very first layer aborted -- on the SACRED gate and equally in the production server. No build flag could avoid it; there was no second path to select. Introduced by 0c2827c (perf(gemma4): dual-GPU FP8 resident, #154), which replaced the portable `vt::GeluAndMul` at this call site with the ROCm-only op in the same commit that added the throw. The fix is the compose the file already uses one function above, in RmsNormPlusAdd: vt::GeluAndMul computes exactly this math, gelu_tanh(gate)*up, but wants the halves adjacent as one [rows, 2D] tensor, so the fallback stages gate and up into a temporary [1, 2n] laid out as [gate | up] and runs the shipped kernel. It therefore matches GeluAndMul BY CONSTRUCTION on every backend that registers it -- CPU, CUDA, Metal, Vulkan -- rather than being a second copy of the gelu math that could drift. ROCm keeps its fused kernel and is untouched. The temporary is read by work queued on `q`, so it is freed only after Synchronize; on CPU that is a no-op, on CUDA it is what stops Free racing the kernel. Freeing on the exception path too. test_ops_activation 21/21 (18/21 with 3 THROWN before this commit) Three cases added, all mutation-checked: swapping gate and up in the staging copies fails every one of them. They pin the golden, byte-identity against GeluAndMul at the real E4B PLE width, and the bf16 dtype the call site actually passes. SCOPE, so this is not read as more than it is: 1. This restores correctness, not speed. The compose costs one allocation, two copies and a synchronize per call, against ROCm's single fused kernel. A real portable kernel (or an elementwise multiply op, which vt does not have -- only MulScalar) would remove all three; that is a perf follow-up, not a correctness one. 2. vt::DualRmsNormPlusRes, vt::MatmulBTAlphaBeta and vt::MatmulBTFp8Channel in this same file still throw by the same pattern. They are NOT reached by the Gemma-4 gate, so they are left alone here and recorded in #377 rather than fixed blind. Refs #377. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude:claude-opus-5 [ClaudeCode]
Summary
Follow-up to merged #140 (ROCm gfx1201 + Gemma-4-26B MoE BF16/FP8).
This PR hardens dual-GPU FP8 expert residency and decode mix so the path fits ~30 GB host RAM and keeps weights on GPU:
MulScalar+Addfor expert weighted sum (default;VT_GEMMA4_HOST_AXPY=1fallback)Lab evidence (2× R9700 gfx1201, ROCm 7.2.x)
<bos>The capital of France is→**Paris**EXIT=0Host OOM root cause was resident upload calling cache-filling dequant for every expert every layer.
Usage
HIP_VISIBLE_DEVICES=0,1 \ VT_GEMMA4_RESIDENT_EXPERTS=1 \ VT_GEMMA4_RESIDENT_GPUS=2 \ ./build-hip/examples/vllm-cli \ --model /path/to/Firworks/gemma-4-26B-A4B-it-fp8 \ --prompt '<bos>The capital of France is' --max-tokens 12 --temperature 0Optional:
VT_GEMMA4_HOST_AXPY=1,VT_GEMMA4_RESIDENT_MAX_LAYERS=NTest plan
Notes